
#include <assert.h>
#include <stdlib.h>
#include <iostream>
#include <iomanip>
#include <string>
#include <math.h>
#include "time.h"
#include "utils.h"
#include "apf.h"
#include "types.h"
#include <stdio.h>
using namespace std;

#define repeat128(S) S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S S

//MAD a=a+b;a=a*b;
__global__ void MAD(unsigned int *ts, unsigned int* out, _DOUBLE_ a1, _DOUBLE_  b1, int its, _DOUBLE_ * result) 
{
  _DOUBLE_ a = a1;
  _DOUBLE_ b = b1;
  _DOUBLE_ c = 0.0;

  unsigned int start_time=0, stop_time=0;
  for (int i=0;i<its;i++) {
    __syncthreads();									
    start_time = clock();							
    repeat128(a=a+b;a=a*b;)
    stop_time = clock();
    c += a + b;	
    printf ("%f\n", c);	
							
  }											
  ts[(blockIdx.x*blockDim.x + threadIdx.x)*2] = start_time;
  ts[(blockIdx.x*blockDim.x + threadIdx.x)*2 +1] = stop_time;	
  result[0] = c;	
		
}
//MAD_BSHARED a=a+b[idx][idy];a=a*b[idx][idy];
__global__ void MAD_BSHARED(unsigned int *ts, unsigned int* out, _DOUBLE_ a1, _DOUBLE_  b1, int its, _DOUBLE_ * result) 
{
  _DOUBLE_ a = a1;
  _DOUBLE_ c = 0.0;

  __shared__ _DOUBLE_ b[BLOCKDIM_X][BLOCKDIM_Y];
  int idx = blockIdx.x;
  int idy = blockIdx.y;
  
  b[idx][idy] = b1;
  __syncthreads();									

  
  unsigned int start_time=0, stop_time=0;
  for (int i=0;i<its;i++) {
    __syncthreads();									
    start_time = clock();							
    repeat128(a=a+b[idx][idy];a=a*b[idx][idy];)
    stop_time = clock();    
    c += a + b[idx][idy];	
    printf ("%f\n", c);	
  }											
  
  ts[(blockIdx.x*blockDim.x + threadIdx.x)*2] = start_time;
  ts[(blockIdx.x*blockDim.x + threadIdx.x)*2 +1] = stop_time;
  result[0] = c;				
}

//MAD2 a=a*b+c
__global__ void MAD2(unsigned int *ts, unsigned int* out, _DOUBLE_ a1, _DOUBLE_  b1,_DOUBLE_  c1, int its, _DOUBLE_ * result) 
{
  _DOUBLE_ a = a1;
  _DOUBLE_ b = b1;
  _DOUBLE_ c = c1;  
  _DOUBLE_ d = 0; 
  unsigned int start_time=0, stop_time=0;
  for (int i=0;i<its;i++) {
    __syncthreads();									
    start_time = clock();							
    repeat128(a=a*b+c;)
    stop_time = clock();
    d += a + b;	
    printf ("%f\n", d);	
					
  }											
  
  ts[(blockIdx.x*blockDim.x + threadIdx.x)*2] = start_time;
  ts[(blockIdx.x*blockDim.x + threadIdx.x)*2 +1] = stop_time;	
  result[0] = d;
}

//MAD2 a=a*b[idx][idy]+c
__global__ void MAD2_BSHARED(unsigned int *ts, unsigned int* out, _DOUBLE_ a1, _DOUBLE_  b1,_DOUBLE_  c1, int its, _DOUBLE_ * result) 
{
  _DOUBLE_ a = a1;
  _DOUBLE_ c = c1; 
  _DOUBLE_ d = 0; 
  __shared__ _DOUBLE_ b[BLOCKDIM_X][BLOCKDIM_Y];
 // __shared__ _DOUBLE_ s[BLOCKDIM_X][BLOCKDIM_Y];

  int idx = blockIdx.x;
  int idy = blockIdx.y;
  
  b[idx][idy] = b1;
  //s[idx][idy] = c1;
  __syncthreads();									
  
  unsigned int start_time=0, stop_time=0;
  for (int i=0;i<its;i++) {
    __syncthreads();									
    start_time = clock();							
    repeat128(a=a*b[idx][idy]+c;)
    stop_time = clock();
    d += a + b[idx][idy]+c;	
    printf ("%f\n", d);								
  }											
  
  ts[(blockIdx.x*blockDim.x + threadIdx.x)*2] = start_time;
  ts[(blockIdx.x*blockDim.x + threadIdx.x)*2 +1] = stop_time;
  result[0] = d;
}

#define _FLOAT_ float

__global__ void SFU(unsigned int *ts, unsigned int* out, _DOUBLE_ a1, _DOUBLE_  b1, int its, _DOUBLE_ * result) 
{
  _FLOAT_ a = 4.0;
  _FLOAT_ b = 4096.0 * 4096 * 4096 * 4096;
  _FLOAT_ c = 0;

  unsigned int start_time=0, stop_time=0;
  do{
    for (int i=0;i<1;i++) {
      __syncthreads();									
      start_time = clock();							
      //repeat128(a=log2(a*a);a=rsqrt(a*a);)
      //repeat128(a=sqrt(b);b=sqrt(a);)
      a = sqrt(b);
      b = sqrt(a);
      a = sqrt(b);
      b = sqrt(a);
      //a = sqrt(b);
      //b = sqrt(a);

    /* a = log2(b);
      b = __log2f(a);
      a = __log2f(b);
      b = __log2f(a);
      a = __log2f(b); 

      b = __log2f(a); */

/*      a = a*b + a;
      b = a*b + b; 
//      b = a + b; b = a*b; 
  //    b = a + b; a = a*b;
      a = a*b + a;
     b = a*b + b;
     //a = a*b + a;
     //b = a*b + b;
            a = a*b + b;
      b = a*b + a;
      a = a*b + b;
      b = a*b + a;
      a = a*b + b; */
      stop_time = clock();

      printf("a: %f\nb: %f\n", a, b);							
      c = a*a + b;


      printf ("c: %f\n", c);								
    }											
  }while(stop_time<start_time);

  ts[(blockIdx.x*blockDim.x + threadIdx.x)*2] = start_time;
  ts[(blockIdx.x*blockDim.x + threadIdx.x)*2 +1] = stop_time;
  result[0] = c;	
}

void measure_MAD(){
  unsigned int timeStamps[2 * 1024];		
  unsigned int *dTimeStamps;
  unsigned int *dOutput;			
  cudaError_t error;
  _DOUBLE_ * result;
  _DOUBLE_ a = 0.0001;
  _DOUBLE_ b = 0.0002;
  _DOUBLE_ c = 0.0004;
  _DOUBLE_ iterations = 2;
  
  //only using one thread??
  dim3 threads(BLOCKDIM_X, BLOCKDIM_Y);
  dim3 grid(BLOCKDIM_X, BLOCKDIM_Y, 1);

  //MAD
  cudaMalloc((void**)&dTimeStamps, sizeof(timeStamps));
  checkCUDAError("Can't allocate timeStamps array");
  cudaMalloc((void**)&dOutput, 4);
  checkCUDAError("Can't allocate dOutput array");
  cudaMalloc((void**)&result, 4);
  checkCUDAError("Can't allocate result");

  MAD <<<grid, threads>>>(dTimeStamps, dOutput, a, b, iterations, result);							
  cudaThreadSynchronize();									
  
  printf("%slatency: ", "MAD a=a+b;a=a*b;\n"); 							
  if ((error = cudaGetLastError()) != cudaSuccess)						
  {												
	  printf("  failed. %s\n\n", cudaGetErrorString(error));					
  }												
  cudaMemcpy(timeStamps, dTimeStamps, sizeof(timeStamps), cudaMemcpyDeviceToHost);					
  cudaThreadSynchronize();									
  printf ("%f cycles\n\n", (timeStamps[1]-timeStamps[0])/(128.0));	

  cudaFree(dTimeStamps);
  checkCUDAError("Can't free timeStamps on device");
  cudaFree(dOutput);
  checkCUDAError("Can't free dOutput on device");
  cudaFree(result);
  checkCUDAError("Can't free result on device");
}

void measure_MAD_BSHARED(){
  unsigned int timeStamps[2 * 1024];		
  unsigned int *dTimeStamps;
  unsigned int *dOutput;			
  cudaError_t error;
  _DOUBLE_ * result;
  _DOUBLE_ a = 0.0001;
  _DOUBLE_ b = 0.0002;
  _DOUBLE_ c = 0.0004;
  _DOUBLE_ iterations = 1;
  
  //only using one thread??
  dim3 threads(BLOCKDIM_X, BLOCKDIM_Y);
  dim3 grid(BLOCKDIM_X, BLOCKDIM_Y, 1);

  //MAD_BSHARED
  cudaMalloc((void**)&dTimeStamps, sizeof(timeStamps));
  checkCUDAError("Can't allocate timeStamps array");
  cudaMalloc((void**)&dOutput, 4);
  checkCUDAError("Can't allocate d_out array");
  cudaMalloc((void**)&result, 4);
  checkCUDAError("Can't allocate result");

  MAD_BSHARED <<<grid, threads>>>(dTimeStamps, dOutput, a, b, iterations,result);							
  cudaThreadSynchronize();									
 
  printf("%slatency: ", "MAD2_BSHARED a=a+b[idx][idy];a=a*b[idx][idy];\n"); 							
  if ((error = cudaGetLastError()) != cudaSuccess)						
  {												
	  printf("  failed. %s\n\n", cudaGetErrorString(error));					
  }												
  cudaMemcpy(timeStamps, dTimeStamps, sizeof(timeStamps), cudaMemcpyDeviceToHost);					
  cudaThreadSynchronize();									
  printf ("%f cycles\n\n", (timeStamps[1]-timeStamps[0]));	

  cudaFree(dTimeStamps);
  checkCUDAError("Can't free timeStamps on device");
  cudaFree(dOutput);
  checkCUDAError("Can't free dOutput on device");
}

//MAD2 a=a*b+c
void measure_MAD2(){
  unsigned int timeStamps[2 * 1024];		
  unsigned int *dTimeStamps;
  unsigned int *dOutput;			
  cudaError_t error;
  _DOUBLE_ * result;
  _DOUBLE_ a = 0.0001;
  _DOUBLE_ b = 0.0002;
  _DOUBLE_ c = 0.0004;
  _DOUBLE_ iterations = 1;
  
  //only using one thread??
  dim3 threads(BLOCKDIM_X, BLOCKDIM_Y);
  dim3 grid(BLOCKDIM_X, BLOCKDIM_Y, 1);

  //MAD2
  cudaMalloc((void**)&dTimeStamps, sizeof(timeStamps));
  checkCUDAError("Can't allocate timeStamps array");
  cudaMalloc((void**)&dOutput, 4);
  checkCUDAError("Can't allocate d_out array");
  cudaMalloc((void**)&result, 4);
  checkCUDAError("Can't allocate result");

  MAD2 <<<grid, threads>>>(dTimeStamps, dOutput, a, b, c, iterations, result);							
  cudaThreadSynchronize();									
 
  printf("%slatency: ", "MAD2 a=a*b+c;\n"); 							
  if ((error = cudaGetLastError()) != cudaSuccess)						
  {												
	  printf("  failed. %s\n\n", cudaGetErrorString(error));					
  }												
  cudaMemcpy(timeStamps, dTimeStamps, sizeof(timeStamps), cudaMemcpyDeviceToHost);					
  cudaThreadSynchronize();									
  printf ("%f cycles\n\n", (timeStamps[1]-timeStamps[0])/128.0);	

  cudaFree(dTimeStamps);
  checkCUDAError("Can't free timeStamps on device");
  cudaFree(dOutput);
  checkCUDAError("Can't free dOutput on device");
  cudaFree(result);
  checkCUDAError("Can't free result on device");
}

void measure_MAD2_BSHARED(){
  unsigned int timeStamps[2 * 1024];		
  unsigned int *dTimeStamps;
  unsigned int *dOutput;			
  cudaError_t error;
  _DOUBLE_ * result;
  _DOUBLE_ a = 0.0001;
  _DOUBLE_ b = 0.0002;
  _DOUBLE_ c = 0.0004;
  _DOUBLE_ iterations = 2;
  
  //only using one thread??
  dim3 threads(BLOCKDIM_X, BLOCKDIM_Y);
  dim3 grid(BLOCKDIM_X, BLOCKDIM_Y, 1);

  //MAD2_BSHARED
  cudaMalloc((void**)&dTimeStamps, sizeof(timeStamps));
  checkCUDAError("Can't allocate timeStamps array");
  cudaMalloc((void**)&dOutput, 4);
  checkCUDAError("Can't allocate d_out array");
  cudaMalloc((void**)&result, 4);
  checkCUDAError("Can't allocate result");

  MAD2_BSHARED <<<grid, threads>>>(dTimeStamps, dOutput, a, b, c, iterations,result);							
  cudaThreadSynchronize();									
 
  printf("%slatency: ", "MAD2_BSHARED a=a*b[idx][idy]+c;\n"); 							
  if ((error = cudaGetLastError()) != cudaSuccess)						
  {												
	  printf("  failed. %s\n\n", cudaGetErrorString(error));					
  }												
  cudaMemcpy(timeStamps, dTimeStamps, sizeof(timeStamps), cudaMemcpyDeviceToHost);					
  cudaThreadSynchronize();									
  printf ("%f cycles\n\n", (timeStamps[1]-timeStamps[0])/128.0);	

  cudaFree(dTimeStamps);
  checkCUDAError("Can't free timeStamps on device");
  cudaFree(dOutput);
  checkCUDAError("Can't free dOutput on device");
  cudaFree(result);
  checkCUDAError("Can't free result on device");
}

void measure_LOG(){
  unsigned int timeStamps[2 * 1024];		
  unsigned int *dTimeStamps;
  unsigned int *dOutput;			
  cudaError_t error;
  _DOUBLE_ * result;
  _DOUBLE_ a = 1024;
  _DOUBLE_ b = 1024;
  _DOUBLE_ c = 1024;
  _DOUBLE_ iterations = 2;
  
  //only using one thread??
  dim3 threads(1, 1);
  dim3 grid(1, 1);

  //LOG_RSQRT
  cudaMalloc((void**)&dTimeStamps, sizeof(timeStamps));
  checkCUDAError("Can't allocate timeStamps array");
  cudaMalloc((void**)&dOutput, 4);
  checkCUDAError("Can't allocate d_out array");
  cudaMalloc((void**)&result, 4);
  checkCUDAError("Can't allocate result");

  SFU <<<1, 1>>>(dTimeStamps, dOutput, a, b, iterations, result);							
  cudaThreadSynchronize();									
 
  printf("%slatency: ", "SFU\n"); 							
  if ((error = cudaGetLastError()) != cudaSuccess)						
  {												
	  printf("  failed. %s\n\n", cudaGetErrorString(error));					
  }												
  cudaMemcpy(timeStamps, dTimeStamps, sizeof(timeStamps), cudaMemcpyDeviceToHost);					
  cudaThreadSynchronize();									
  printf ("%i cycles\n\n", (timeStamps[1]-timeStamps[0]));	

  cudaFree(dTimeStamps);
  checkCUDAError("Can't free dTimeStamps on device");
  cudaFree(dOutput);
  checkCUDAError("Can't free dOutput on device");
  cudaFree(result);
  checkCUDAError("Can't free result on device");
  
}

void measure_latency(){
  //measure_MAD();
  //measure_MAD_BSHARED();
  //measure_MAD2();
  //measure_MAD2_BSHARED();
  measure_LOG();
}